[BugFix][DP/EP] Fix CUTLASS MLA hang under load #26026

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Merged

vllm-bot merged 2 commits into vllm-project:main from neuralmagic:lwilkinson/fix-cutlass-mla-hang

Oct 1, 2025

+32 −32

Collaborator

LucasWilkinson commented Oct 1, 2025 •

edited by github-actions bot

Loading

The early return in compute( calls arrive:

      cutlass::arch::NamedBarrier(
          (kNumComputeWarps + kNumLoadWarps) * NumThreadsPerWarp,
          kNamedBarrierEpilogue
      ).arrive();

but didn't have any barrier before looping around and calling it again causing a deadlock when the load warps waits on:

cutlass::arch::NamedBarrier((kNumComputeWarps + kNumLoadWarps) * NumThreadsPerWarp, kNamedBarrierEpilogue).arrive_and_wait();

Co-authored-by: Robert Shaw robshaw@redhat.com

gemini-code-assist bot reviewed

View reviewed changes

Contributor

gemini-code-assist bot left a comment

Code Review

This pull request addresses a potential hang in the CUTLASS MLA kernel under load. The changes are mostly stylistic, involving code re-indentation. However, there is one critical bug fix that changes a cutlass::arch::NamedBarrier::arrive() call to arrive_and_wait(). This change correctly resolves a race condition that could lead to a deadlock between compute and load warps, which is the likely cause of the hang. My review confirms this fix is correct and critical.

csrc/attention/mla/cutlass_sm100_mla/kernel/sm100_fmha_mla_tma_warpspecialized.hpp Show resolved Hide resolved

LucasWilkinson requested a review from robertgshaw2-redhat

October 1, 2025 14:07

robertgshaw2-redhat added this to the v0.11.0 Cherry Picks milestone

Collaborator

robertgshaw2-redhat commented Oct 1, 2025

TODOs (follow-up)

investigate whether this can solve the hangs for num_kv_splits>1
investigate a dynamic scheduler

LucasWilkinson force-pushed the lwilkinson/fix-cutlass-mla-hang branch from 746ef1a to a36b036 Compare

October 1, 2025 14:10


          arrive and wait

3a8a634

Co-authored-by: Robert Shaw <robshaw@redhat.com>
Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>

LucasWilkinson force-pushed the lwilkinson/fix-cutlass-mla-hang branch from a36b036 to 3a8a634 Compare

October 1, 2025 14:12

mgoin approved these changes

View reviewed changes

mgoin added bug ready deepseek labels

robertgshaw2-redhat reviewed

View reviewed changes

csrc/attention/mla/cutlass_sm100_mla/kernel/sm100_fmha_mla_tma_warpspecialized.hpp

    
                        (kNumComputeWarps + kNumLoadWarps) * NumThreadsPerWarp,

                        kNamedBarrierEpilogue

                    ).arrive();

                    ).arrive_and_wait();

Collaborator

robertgshaw2-redhat Oct 1, 2025

note for reviewer --- this is the line that fixes the deadlock

robertgshaw2-redhat changed the title ~~[BugFix] Fix CUTLASS MLA hang under load~~ [BugFix][DP/EP] Fix CUTLASS MLA hang under load


          remove tabs

3c496cb

Co-authored-by: rshaw@neuralmagic.com <rshaw@neuralmagic.com>
Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>

tlrmchlsmth approved these changes

View reviewed changes

LucasWilkinson mentioned this pull request

revert max split heuristics neuralmagic/vllm#121

Draft

tlrmchlsmth mentioned this pull request

[Bug][WideEP]: API server hangs during DP=16 + DBO on B200 at high concurrency #25491

Closed

1 task

Member

mgoin commented Oct 1, 2025 •

edited

Loading

Failing tests are known on main and being resolved. Blackwell tests are green, merging

vllm-bot merged commit 1726e93 into vllm-project:main

79 of 84 checks passed

soldni pushed a commit to soldni/vllm that referenced this pull request


          [BugFix][DP/EP] Fix CUTLASS MLA hang under load (vllm-project#26026)

bf632d5

Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
Co-authored-by: Robert Shaw <robshaw@redhat.com>
Co-authored-by: rshaw@neuralmagic.com <rshaw@neuralmagic.com>

simon-mo pushed a commit that referenced this pull request


          [BugFix][DP/EP] Fix CUTLASS MLA hang under load (#26026)

ebce361

Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
Co-authored-by: Robert Shaw <robshaw@redhat.com>
Co-authored-by: rshaw@neuralmagic.com <rshaw@neuralmagic.com>
Signed-off-by: simon-mo <simon.mo@hey.com>

pdasigi pushed a commit to pdasigi/vllm that referenced this pull request


          [BugFix][DP/EP] Fix CUTLASS MLA hang under load (vllm-project#26026)

30ea5a5

Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
Co-authored-by: Robert Shaw <robshaw@redhat.com>
Co-authored-by: rshaw@neuralmagic.com <rshaw@neuralmagic.com>

yewentao256 pushed a commit that referenced this pull request


          [BugFix][DP/EP] Fix CUTLASS MLA hang under load (#26026)

90529ce

Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
Co-authored-by: Robert Shaw <robshaw@redhat.com>
Co-authored-by: rshaw@neuralmagic.com <rshaw@neuralmagic.com>
Signed-off-by: yewentao256 <zhyanwentao@126.com>

tomeras91 pushed a commit to tomeras91/vllm that referenced this pull request


          [BugFix][DP/EP] Fix CUTLASS MLA hang under load (vllm-project#26026)

f882803

Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
Co-authored-by: Robert Shaw <robshaw@redhat.com>
Co-authored-by: rshaw@neuralmagic.com <rshaw@neuralmagic.com>
Signed-off-by: Tomer Asida <57313761+tomeras91@users.noreply.github.com>

southfreebird pushed a commit to southfreebird/vllm that referenced this pull request


          [BugFix][DP/EP] Fix CUTLASS MLA hang under load (vllm-project#26026)

9b59b37

Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
Co-authored-by: Robert Shaw <robshaw@redhat.com>
Co-authored-by: rshaw@neuralmagic.com <rshaw@neuralmagic.com>

xuebwang-amd pushed a commit to xuebwang-amd/vllm that referenced this pull request


          [BugFix][DP/EP] Fix CUTLASS MLA hang under load (vllm-project#26026)

b7b276c

Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
Co-authored-by: Robert Shaw <robshaw@redhat.com>
Co-authored-by: rshaw@neuralmagic.com <rshaw@neuralmagic.com>
Signed-off-by: xuebwang-amd <xuebwang@amd.com>

choprahetarth pushed a commit to Tandemn-Labs/vllm that referenced this pull request


          [BugFix][DP/EP] Fix CUTLASS MLA hang under load (vllm-project#26026)

3c97124

Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
Co-authored-by: Robert Shaw <robshaw@redhat.com>
Co-authored-by: rshaw@neuralmagic.com <rshaw@neuralmagic.com>
Signed-off-by: simon-mo <simon.mo@hey.com>

shyeh25 pushed a commit to shyeh25/vllm that referenced this pull request


          Fix CUTLASS MLA hang under load (vllm-project#26026)

89581c5

Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
Co-authored-by: Robert Shaw <robshaw@redhat.com>
Co-authored-by: rshaw@neuralmagic.com <rshaw@neuralmagic.com>
Signed-off-by: simon-mo <simon.mo@hey.com>

MatthewBonanni mentioned this pull request

[Attention] Tune CUTLASS MLA num_splits #26846

Merged

5 tasks

lywa1998 pushed a commit to lywa1998/vllm that referenced this pull request


          [BugFix][DP/EP] Fix CUTLASS MLA hang under load (vllm-project#26026)

27a2ca2

Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
Co-authored-by: Robert Shaw <robshaw@redhat.com>
Co-authored-by: rshaw@neuralmagic.com <rshaw@neuralmagic.com>

alhridoy pushed a commit to alhridoy/vllm that referenced this pull request


          [BugFix][DP/EP] Fix CUTLASS MLA hang under load (vllm-project#26026)

ff95807

Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
Co-authored-by: Robert Shaw <robshaw@redhat.com>
Co-authored-by: rshaw@neuralmagic.com <rshaw@neuralmagic.com>

xuebwang-amd pushed a commit to xuebwang-amd/vllm that referenced this pull request


          [BugFix][DP/EP] Fix CUTLASS MLA hang under load (vllm-project#26026)

e7180ee

Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
Co-authored-by: Robert Shaw <robshaw@redhat.com>
Co-authored-by: rshaw@neuralmagic.com <rshaw@neuralmagic.com>
Signed-off-by: xuebwang-amd <xuebwang@amd.com>

rtourgeman pushed a commit to rtourgeman/vllm that referenced this pull request


          [BugFix][DP/EP] Fix CUTLASS MLA hang under load (vllm-project#26026)

a83a90f

Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
Co-authored-by: Robert Shaw <robshaw@redhat.com>
Co-authored-by: rshaw@neuralmagic.com <rshaw@neuralmagic.com>

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

bug deepseek ready